/*
*  The data needed for ray-intersection detection	
*  Cache friendly
*/
struct TriAccel
{
    /* first 16 byte half cache line */
    /* plane */
    float nu;       /* normal.u / normal.k */ 
    float nv;       /* normal.v / normla.k */
    float nd;       /* constant value in plane equation */
    int k;          /* projection dimension */

    /* second 16 byte half cache line */
    /* line equation for line ac */
    float b_nu;     
    float b_nv;
    float b_d;
    int pad_1;

    /* second 16 byte half cache line */
    /* line equation for line ab */
    float c_nu; 
    float c_nv;
    float c_d;
    unsigned int id;
    // int pad_2;
};

struct Triangle 
{
	unsigned int vertIndex[3];
	unsigned int id; 
};

/* Vertex Buffer
*/
struct VertexData
{
    float3    *posArray;
    float3    *normalArray;
    float3    *texcoordArray;
    unsigned int numVerts;
	
	VertexData()
		: posArray(NULL), normalArray(NULL), texcoordArray(NULL), numVerts(NULL)
	{}
};

/* Index Buffer
*/ 
struct IndexData
{
    unsigned int *indices;
    unsigned int numFaces;
	
	IndexData() 
		: indices(NULL), numFaces(0)
	{}
};

/* Triangle Buffer 
*/
struct TriangleData
{
    Triangle    *triangles;
    unsigned int numTri;
	
	TriangleData()
		: triangles(NULL), numTri(0)
	{}
};

__device__ void 